#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0

__global__ void standard_kernel(float inputValue, float *res)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    res[idx] = powf(inputValue, 2.0f) * idx;
}

__global__ void intrinsic_kernel(float inputValue, float *res)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    res[idx] = __powf(inputValue, 2.0f) * idx;
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);

    int nElem = 1 << 15;
    size_t nBytes = nElem * sizeof(float);

    float input_value = 10.0f;
    float *host_array_out = (float *)malloc(nBytes);
    memset(host_array_out, 0, nBytes);

    float *device_array_out = nullptr;
    CHECK(cudaMalloc((float **)&device_array_out, nBytes));
    CHECK(cudaMemcpy(device_array_out, host_array_out, nBytes, cudaMemcpyHostToDevice));

    dim3 block_size = block_dims(nElem);
    dim3 grid_size = grid_dims(nElem);

    standard_kernel<<<grid_size, block_size>>>(input_value, device_array_out);
    CHECK(cudaDeviceSynchronize());

    intrinsic_kernel<<<grid_size, block_size>>>(input_value, device_array_out);
    CHECK(cudaDeviceSynchronize());

    free(host_array_out);
    cudaFree(device_array_out);

    return 0;
}